//-------------------------------------------------------------------
//	Copyright (c) 2013  Zhirnov Andrey
//	This file is part of the "Noise Library" project.
//	See copyright notice in "linNoiseCL.h".
//-------------------------------------------------------------------

#include "CProgramCL.h"
#include "../Context/CContext.h"
#include "../Source/Controls/CSourceOutput.h"

namespace Noise
{

	//
	// Profiler
	//

	struct Profiler
	{
		string	_name;
		double	_timeStart;

		Profiler ()
		{
			Start();
		}

		Profiler (const char *name) : _name(name)
		{
			Start();
		}

		~Profiler ()
		{
			Stop();
		}

#ifndef _NOISE_CROSS_PLATFORM
		void Start ()
		{
			_timeStart = CONTEXT()->GetTime();
		}

		void Stop () const
		{
			LOG( (string(_name) << (CONTEXT()->GetTime() - _timeStart) << " s.").cstr(), e_log_msg::DEBUG );
		}
#else
		void Start () {}
		void Stop () const {}
#endif
	};


	// constructor
	CProgramCL::CProgramCL ()
	{
	}

	// destructor
	CProgramCL::~CProgramCL ()
	{
	}

	// Delete
	void CProgramCL::Delete ()
	{
		CProgram::Delete();
		
		_format = EPixelFormat(0);

		_prog.Destroy();
		_kernel.Destroy();
		_compileLog.Clear();
	}

	// _GetOptions
	bool CProgramCL::_GetOptions (EBuildOptions options, cl_prog_options::type &eOpFlags, string &sOptions) const
	{
		// TODO
		return true;
	}

	// _GetPreProcessor
	void CProgramCL::_AddPreProcessor (EBuildOptions options)
	{
		const int	op_mask		= BUILD_FLOAT16 | BUILD_FLOAT32 | BUILD_FLOAT64 |
								  BUILD_OUTPUT_IMAGE_2D | BUILD_OUTPUT_IMAGE_3D;
		int			new_options = 0;

		if ( EnumCmp( options, BUILD_FLOAT16 ) )
		{
			_srcData.source
				<< "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n"
				<< "#define real			half\n"
				<< "#define real2			half2\n"
				<< "#define real3			half3\n"
				<< "#define real4			half4\n"
				<< "#define real8			half8\n"
				<< "#define real16			half16\n\n"
				<< "#define convert_real2	convert_half2\n"
				<< "#define convert_real3	convert_half3\n"
				<< "#define convert_real4	convert_half4\n"
				<< "#define convert_real8	convert_half8\n"
				<< "#define convert_real16	convert_half16\n\n";
			new_options |= BUILD_FLOAT16;
		}
		else
		if ( EnumCmp( options, BUILD_FLOAT32 ) )
		{
			_srcData.source
				<< "#define real			float\n"
				<< "#define real2			float2\n"
				<< "#define real3			float3\n"
				<< "#define real4			float4\n"
				<< "#define real8			float8\n"
				<< "#define real16			float16\n\n"
				<< "#define convert_real2	convert_float2\n"
				<< "#define convert_real3	convert_float3\n"
				<< "#define convert_real4	convert_float4\n"
				<< "#define convert_real8	convert_float8\n"
				<< "#define convert_real16	convert_float16\n\n";
			new_options |= BUILD_FLOAT32;
		}
		else
		if ( EnumCmp( options, BUILD_FLOAT64 ) )
		{
			_srcData.source
				<< "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n"
				<< "#define real			double\n"
				<< "#define real2			double2\n"
				<< "#define real3			double3\n"
				<< "#define real4			double4\n"
				<< "#define real8			double8\n"
				<< "#define real16			double16\n\n"
				<< "#define convert_real2	convert_double2\n"
				<< "#define convert_real3	convert_double3\n"
				<< "#define convert_real4	convert_double4\n"
				<< "#define convert_real8	convert_double8\n"
				<< "#define convert_real16	convert_double16\n\n";
			new_options |= BUILD_FLOAT64;
		}

		if ( EnumCmp( options, BUILD_OUTPUT_SIGNED_INT ) )
		{
			_srcData.source << "#define __imageWrite	write_imagei\n\n";
			new_options |= BUILD_OUTPUT_SIGNED_INT;
		}
		else
		if ( EnumCmp( options, BUILD_OUTPUT_UNSIGNED_INT ) )
		{
			_srcData.source << "#define __imageWrite	write_imageui\n\n";
			new_options |= BUILD_OUTPUT_UNSIGNED_INT;
		}
		else
		{
			_srcData.source << "#define __imageWrite	write_imagef\n\n";
			new_options |= BUILD_OUTPUT_FLOAT;
		}

		if ( EnumCmp( options, BUILD_OUTPUT_IMAGE_2D ) )
		{
			_srcData.source
				<< "#define outputImage_t	image2d_t\n"
				<< "#define writeImage( _img, _coord, _value )	__imageWrite( _img, _coord.xy, _value )\n\n";
			new_options |= BUILD_OUTPUT_IMAGE_2D;
		}
		else
		if ( EnumCmp( options, BUILD_OUTPUT_IMAGE_3D ) )
		{
			_srcData.source
				<< "#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable\n"
				<< "#define outputImage_t	image3d_t\n"
				<< "#define writeImage( _img, _coord, _value )	__imageWrite( _img, _coord.xyz, _value )\n\n";
			new_options |= BUILD_OUTPUT_IMAGE_3D;
		}

		_srcData.source
			<< "#define in \n"
			<< "#define static_const	__constant\n"
			<< "#define globalSync()	barrier(CLK_GLOBAL_MEM_FENCE)\n"
			<< "#define localSync()		barrier(CLK_LOCAL_MEM_FENCE)\n\n";

		_options = EBuildOptions( (_options & ~op_mask) | new_options );
	}
	
	// _GetProgramMainSrc
	void CProgramCL::_AddProgramMainSrc (CSourcePtr source)
	{
		_srcData.source
			<< "\n\n"
			<< "real3 NTransformCoord (in int3 coord, in int3 size, in int3 offset, in float3 scale, in float3 bias)\n"
			<< "{\n"
			<< "	return convert_real3(coord + offset) / convert_real3(size) * convert_real3(scale) + convert_real3(bias);\n"
			<< "}\n\n"

			<< "__kernel void libNoiseEntry ( write_only outputImage_t outImage, float4 scale, float4 bias ) \n"
			<< "{\n"
			<< "	int3	size  = (int3)( get_global_size(0),   get_global_size(1),   get_global_size(2)   );\n"
			<< "	int3	off   = (int3)( get_global_offset(0), get_global_offset(1), get_global_offset(2) );\n"
			<< "	int3	id    = (int3)( get_global_id(0),     get_global_id(1),     get_global_id(2)     );\n"
			<< "	real3	pos   = NTransformCoord( id, size, off, scale.xyz, bias.xyz );\n"
			<< "	int3	coord = id + off;\n"
			<< "	real4	res   = (real4)( " << CSource::FunctionName() << source->GetSourceId() << "( pos ) );\n"
			<< "	writeImage( outImage, coord, res );\n"
			<< "}\n";
	}
	
	// _CreateDefaultOutput
	void CProgramCL::_CreateDefaultOutput (CSourcePtr source, CSourcePtr &outSource)
	{
		if ( not CSource::HasColorOutput( source ) )
		{
			outSource = CONTEXT()->CreateSource( RENDER_OUTPUT );
			outSource->SetArg( SOURCE_0, source );
		}
		else
			outSource = source;
		
		_format = outSource.To<ISourceOutputPtr>()->GetFormat();
	}

	// Build
	EResult CProgramCL::Build (CSourcePtr source, EBuildOptions options)
	{
		_srcData.Clear();

		CSourcePtr	out_src;

		_AddPreProcessor( options );
		_CreateDefaultOutput( source, out_src );

		CHECK_RETURN_R( CProgram::Build( out_src, options ) == OK, INVALID_OPERATION );

		string					s_macro;
		cl_prog_options::type	e_flags = cl_prog_options::CL_1_1;

		_AddProgramMainSrc( out_src );

		CHECK_RETURN_R( _GetOptions( options, e_flags, s_macro ), INVALID_OPERATION );
		CHECK_RETURN_R( _prog.Create( _srcData.source.cstr() ), INVALID_OPERATION );

		LOG( "Program source:", e_log_msg::DEBUG );
		LOG( _srcData.source.cstr(), e_log_msg::INFO );

		Profiler	prof("Compilation time: ");

		if ( not _prog.Build( e_flags, s_macro ) )
		{
			_prog.GetBuildLog( _compileLog );
			LOG( _compileLog.cstr(), e_log_msg::ERROR_NL_1 );
			return INVALID_OPERATION;
		}

		//CHECK_WARNING( _prog.GetBinary( _binarySrc ) );

		CHECK_RETURN_R( _kernel.Create( &_prog, "libNoiseEntry" ), INVALID_OPERATION );
		return OK;
	}

	// GetInteger
	EResult CProgramCL::GetInteger (EProgramQuery pname, int *v, int count) const
	{
		CHECK_ARGUMENT_R( count >= 1, INVALID_VALUE );

		switch ( pname )
		{
			case PROGRAM_BUILD_STATE :
				v[0] = _prog.IsBuilded();
				return OK;

			case PROGRAM_TYPE :
				v[0] = OPENCL_PROGRAM;
				return OK;

			case PROGRAM_OUT_FORMAT :
				v[0] = _format;
				return OK;

			case PROGRAM_BUILD_LOG_LENGTH :
				v[0] = _srcData.buildLog.Size();
				return OK;

			case PROGRAM_COMPILE_LOG_LENGTH :
				v[0] = _compileLog.Size();
				return OK;

			case PROGRAM_SOURCE_LENGTH :
				v[0] = _srcData.source.Size();
				return OK;

			case PROGRAM_BINARY_LENGTH :
				v[0] = _binarySrc.Size();
				return OK;

			case PROGRAM_CL_POINTER :
				CHECK_ARGUMENT_R( ((uint)count >= ArraySize< cl_program, int >::value), INVALID_VALUE );
				*(cl_program *)v = _prog.Id();
				return OK;

			case PROGRAM_BUILD_OPTIONS :
				v[0] = _options;
				return OK;

			case PROGRAM_MAX_IMAGE_SIZE :
				MemCopy( v, CONTEXT()->GetMaxImageSize().ptr(), Clamp( count * sizeof(int), sizeof(uint), sizeof(uvec3) ) );
				return OK;

			case PROGRAM_NUM_IMAGE_FORMATS :
				v[0] = GetSupportedFormats().Count();
				return OK;

			case PROGRAM_IMAGE_FORMATS :
				CHECK_ARGUMENT_R( (uint)count >= GetSupportedFormats().Count(), INVALID_VALUE );
				MemCopy( v, GetSupportedFormats().ptr(), GetSupportedFormats().Size() );
				return OK;
		}

		return CProgram::GetInteger( pname, v, count );
	}

	// GetString
	EResult CProgramCL::GetString (EProgramQuery pname, char *buf, int length) const
	{
		switch ( pname )
		{
			case PROGRAM_BUILD_LOG :
				CHECK_ARGUMENT_R( (uint)length >= _srcData.buildLog.Size(), INVALID_VALUE );
				MemCopy( buf, _srcData.buildLog.cstr(), _srcData.buildLog.Size() );
				return OK;

			case PROGRAM_COMPILE_LOG :
				CHECK_ARGUMENT_R( (uint)length >= _compileLog.Size(), INVALID_VALUE );
				MemCopy( buf, _compileLog.cstr(), _compileLog.Size() );
				return OK;

			case PROGRAM_SOURCE :
				CHECK_ARGUMENT_R( (uint)length >= _srcData.source.Size(), INVALID_VALUE );
				MemCopy( buf, _srcData.source.cstr(), _srcData.source.Size() );
				return OK;

			case PROGRAM_BINARY :
				CHECK_ARGUMENT_R( (uint)length >= _binarySrc.Size(), INVALID_VALUE );
				MemCopy( buf, _binarySrc.raw_ptr(), _binarySrc.Size() );
				return OK;
		}
		return INVALID_ENUM;
	}
	
	// _RenderImageCL
	bool CProgramCL::_RenderImageCL (cl_mem imageId, const svec3 &offset, const svec3 &dim)
	{
		_kernel.SetArg( 0, imageId );
		_kernel.SetArg( 1, fvec4(_coordScale, 0.0f) );
		_kernel.SetArg( 2, fvec4(_coordBias,  0.0f) );

		CHECK_RETURN( _kernel.EnqueueNDRange( CONTEXT()->GetQueue(), offset.To<svec3>(), dim.To<svec3>(), svec3(0) ) );
		return true;
	}

	// RenderToMemory
	EResult CProgramCL::RenderToMemory (const uvec3 &dim, const uvec2 &sizes, EPixelFormat format, EPixelType type, void *pixels)
	{
		CHECK_ARGUMENT_R( pixels != null, INVALID_VALUE );

		real3			scale;
		CLImage			image;
		cl_img_format	fmt;
		cl_memory::type	flags = cl_memory::WRITE_ONLY;

		switch ( format )
		{
			case R		: fmt.eOrder = cl_img_channel_order::R;		break;
			case RG		: fmt.eOrder = cl_img_channel_order::RG;	break;
			case RGB	: fmt.eOrder = cl_img_channel_order::RGB;	break;
			case RGBA	: fmt.eOrder = cl_img_channel_order::RGBA;	break;
		}

		if ( EnumCmp( _options, BUILD_OUTPUT_SIGNED_INT | BUILD_OUTPUT_UNSIGNED_INT ) )
		{
			switch ( type )
			{
				case SINT8			: fmt.eChannel = cl_img_channel::SINT8;				break;
				case SINT16			: fmt.eChannel = cl_img_channel::SINT16;			break;
				case SINT32			: fmt.eChannel = cl_img_channel::SINT32;			break;
				case UINT8			: fmt.eChannel = cl_img_channel::UINT8;				break;
				case UINT16			: fmt.eChannel = cl_img_channel::UINT16;			break;
				case UINT32			: fmt.eChannel = cl_img_channel::UINT32;			break;
			}
		}
		else
		//if ( EnumCmp( _options, BUILD_OUTPUT_FLOAT ) )
		{
			switch ( type )
			{
				case FLOAT16		: fmt.eChannel = cl_img_channel::HALF;				break;
				case FLOAT32		: fmt.eChannel = cl_img_channel::FLOAT;				break;
				case SNORM_8		: fmt.eChannel = cl_img_channel::SNORM_INT8;		break;
				case SNORM_16		: fmt.eChannel = cl_img_channel::SNORM_INT16;		break;
				case UNORM_8		: fmt.eChannel = cl_img_channel::UNORM_INT8;		break;
				case UNORM_16		: fmt.eChannel = cl_img_channel::UNORM_INT16;		break;
				case UNORM_5_6_5	: fmt.eChannel = cl_img_channel::UNORM_SHORT_565;	fmt.eOrder = cl_img_channel_order::RGB;	break;
				case UNORM_10_10_10	: fmt.eChannel = cl_img_channel::UNORM_INT_101010;	fmt.eOrder = cl_img_channel_order::RGB;	break;
			}
			/*if ( type == FLOAT64_UINT32X2 and (format == R or format == RG) )
			{
				fmt.eChannel = cl_img_channel::UINT32;
				fmt.eOrder	 = format == R ? cl_img_channel_order::RG : cl_img_channel_order::RGBA;
			}*/
		}

		{
			Profiler	perf("Rendering time: ");

			CHECK_VALUE_R( fmt.eChannel != cl_img_channel::type(0) or fmt.eOrder == cl_img_channel_order::type(0), INVALID_ENUM );
			CHECK_RETURN_R( image.Create2D( dim.xy(), flags, fmt, pixels, sizes.x ), INVALID_OPERATION );
		
			CHECK_RETURN_R( _RenderImageCL( image.Id(), svec3(0), dim ), INVALID_OPERATION );

			CHECK_RETURN_R( image.EnqueueRead( CONTEXT()->GetQueue(), false, svec3(0),
								svec3(dim), svec2(sizes), pixels ), INVALID_OPERATION );

			CONTEXT()->GetQueue()->Finish();
		}

		if ( _enableViewer )
		{
			CONTEXT()->GetImageViewer()->DrawImage( dim.xy(), format, type, pixels );
		}

		return OK;
	}

	// RenderToImage
	EResult CProgramCL::RenderToImageCL (cl_mem imageId)
	{
		CHECK_ARGUMENT_R( imageId != null, INVALID_VALUE );

		svec3	dim;
		cl_int	i_cl_err = 0;

		Profiler	prof("Rendering time: ");

		CL_CALL( i_cl_err = clGetImageInfo( imageId, CL_IMAGE_WIDTH,  sizeof(dim.x), &dim.x, null ) );
		CL_CALL( i_cl_err = clGetImageInfo( imageId, CL_IMAGE_HEIGHT, sizeof(dim.y), &dim.y, null ) );
		CL_CALL( i_cl_err = clGetImageInfo( imageId, CL_IMAGE_DEPTH,  sizeof(dim.z), &dim.z, null ) );

		CHECK_RETURN_R( _RenderImageCL( imageId, svec3(0), dim ), INVALID_OPERATION );
		return OK;
	}
	
	// RenderToSubImageCL
	EResult CProgramCL::RenderToSubImageCL (cl_mem imageId, const svec3 &offset, const svec3 &dim)
	{
		CHECK_ARGUMENT_R( imageId != null, INVALID_VALUE );
		
		Profiler	prof("Rendering time: ");

		CHECK_RETURN_R( _RenderImageCL( imageId, offset, dim ), INVALID_OPERATION );
		return OK;
	}

#ifndef _NOISE_CROSS_PLATFORM
	// GetBindedTexture
	void GetBindedTexture (GLenum &target, GLuint &id)
	{
		static const GLenum	targets[] = {
			GL_TEXTURE_BINDING_1D,						GL_TEXTURE_1D,
			GL_TEXTURE_BINDING_1D_ARRAY,				GL_TEXTURE_1D_ARRAY,
			GL_TEXTURE_BINDING_2D,						GL_TEXTURE_2D,
			GL_TEXTURE_BINDING_2D_ARRAY,				GL_TEXTURE_2D_ARRAY,
			GL_TEXTURE_BINDING_2D_MULTISAMPLE,			GL_TEXTURE_2D_MULTISAMPLE,
			GL_TEXTURE_BINDING_2D_MULTISAMPLE_ARRAY,	GL_TEXTURE_2D_MULTISAMPLE_ARRAY,
			GL_TEXTURE_BINDING_3D,						GL_TEXTURE_3D,
			GL_TEXTURE_BINDING_BUFFER,					GL_TEXTURE_BUFFER,
			GL_TEXTURE_BINDING_CUBE_MAP,				GL_TEXTURE_CUBE_MAP,
			GL_TEXTURE_BINDING_RECTANGLE,				GL_TEXTURE_RECTANGLE
		};

		for (uint i = 0; i < count_of(targets)>>1; ++i)
		{
			GL_CALL( glGetIntegerv( targets[i<<1], (GLint *) &id ) );

			if ( id != 0 )
			{
				target = targets[(i<<1)+1];
				return;
			}
		}
		target = GL_TEXTURE_2D;
		id = 0;
	}

	// RenderToTextureGL
	EResult CProgramCL::RenderToTextureGL (GLenum target, GLuint id, int level)
	{
		CLImage			image;
		cl_memory::type	flags = cl_memory::WRITE_ONLY;

		GLenum			prev_target = 0;
		GLuint			prev_id = 0;
		ivec3			dim;
		

		OPENGL_CLEAR_ERRORS();

		GetBindedTexture( prev_target, prev_id );

		GL_CALL( glBindTexture( target, id ) );
		GL_CALL( glGetTexLevelParameteriv( target, level, GL_TEXTURE_WIDTH,  &dim.x ) );
		GL_CALL( glGetTexLevelParameteriv( target, level, GL_TEXTURE_HEIGHT, &dim.y ) );
		GL_CALL( glGetTexLevelParameteriv( target, level, GL_TEXTURE_DEPTH,  &dim.z ) );

		EResult res = INVALID_OPERATION;
		{
			Profiler	prof("Rendering time: ");

			if ( image.CreateGL( target, id, flags, level ) )
			{
				cl_mem imageId = image.Id();

				res = CONTEXT()->GetQueue()->EnqueueAcquireGLObjects( clmem_uni_array( &imageId, 1 ) ) and
					 _RenderImageCL( image.Id(), svec3(0), dim.To<svec3>() ) and
					 CONTEXT()->GetQueue()->EnqueueReleaseGLObjects( clmem_uni_array( &imageId, 1 ) ) ?
					 OK : INVALID_OPERATION;
			}
		}

		GL_CALL( glBindTexture( prev_target, prev_id ) );

		if ( res == OK and _enableViewer )
		{
			CONTEXT()->GetImageViewer()->DrawTexture( id );
		}

		return res;
	}

#else

	EResult CProgramCL::RenderToTextureGL (GLenum target, GLuint id, int level)
	{
		return _NO_IMPLEMENTATION;
	}

#endif	// _NOISE_CROSS_PLATFORM

	// RenderToTextureDX
	EResult CProgramCL::RenderToTextureDX (void *dxTex)
	{
		return _NO_IMPLEMENTATION;
	}

	// CreateCLProgram
	bool CProgramCL::CreateProgram (CProgramPtr &prog, CContext *context, EProgram type)
	{
		ASSERT( type == OPENCL_PROGRAM );

		prog = new CProgramCL();
		return true;
	}

	// GetSupportedFormats
	const uni_array<const CProgram::Format> CProgramCL::GetSupportedFormats ()
	{
		static const Format	_formats[] = {
			//{ R,	SINT8 },
			//{ R,	SINT16 },
			//{ R,	SINT32 },
			//{ R,	UINT8 },
			//{ R,	UINT16 },
			//{ R,	UINT32 },
			{ R,	FLOAT16 },	{ RG,	FLOAT16 },	{ RGB,	FLOAT16 },	{ RGBA,	FLOAT16 },
			{ R,	FLOAT32 },	{ RG,	FLOAT32 },	{ RGB,	FLOAT32 },	{ RGBA,	FLOAT32 },
			//{ R,	FLOAT64_UINT32X2 },				{ RG,	FLOAT64_UINT32X2 },
			//{ R,	FLOAT64 },
			{ R,	SNORM_8 },	{ RG,	SNORM_8 },	{ RGB,	SNORM_8 },	{ RGBA,	SNORM_8 },
			{ R,	SNORM_16 },	{ RG,	SNORM_16 },	{ RGB,	SNORM_16 },	{ RGBA,	SNORM_16 },
			{ R,	UNORM_8 },	{ RG,	UNORM_8 },	{ RGB,	UNORM_8 },	{ RGBA,	UNORM_8 },
			{ R,	UNORM_16 },	{ RG,	UNORM_16 },	{ RGB,	UNORM_16 },	{ RGBA,	UNORM_16 },
			{ RGB,	UNORM_5_6_5 },
			{ RGB,	UNORM_10_10_10 }
		};

		return uni_array<const Format>( _formats );
	}

}	// Noise